


__kernel void f_aiInputLayerFeedforward(__global float* inData, __global half* WeightsNode, __global half* nodeVar, uint NodeNum) {
    size_t id = get_group_id(0);
    
    nodeVar[id] = 1.0 / exp(inData[id] * WeightsNode[id]);
    
    if(id == 0){
        size_t work_size[2];
        work_size[0] = sqrt(NodeNum)+1;
        work_size[1] = work_size[0];
        ndrange_t ndrange1 = ndrange_2D(work_size);
        void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet(Hash, rangeOffset, rangeNum, eNum, Step); };
        int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
        
    }
    
    
}

//卷积
void f_ConvolutionNet(__global half* nodeVar, __global half* WeightsNode, __global half* newNodeVar, uint* weightsOffset, ushort layer) {
    
    size_t id = get_group_id(0);
    
    __local float gnewNodeVar = 0.0;
    
#ifdef DEF_2_GAP
    half2 v1 = vload2(id, nodeVar);
    v1 *= WeightsNode[id];
    atomiv_add(gnewNodeVar, v1.x + v1.y);
    
#elif DEF_3_GAP
    half3 v1 = vload3(id, nodeVar);
    v1 *= WeightsNode[id];
    atomiv_add(gnewNodeVar, v1.x + v1.y + v1.z);
    
#elif DEF_4_GAP
    half4 v1 = vload4(id, nodeVar);
    v1 *= WeightsNode[id];
    atomiv_add(gnewNodeVar, v1.s0 + v1.s1 + v1.s2 + v1.s3);
    
#elif DEF_8_GAP
    half8 v1 = vload8(id, nodeVar);
    v1 *= WeightsNode[id];
    atomiv_add(gnewNodeVar, v1.s0 + v1.s1 + v1.s2 + v1.s3 + v1.s4 + v1.s5 + v1.s6 + v1.s7);
    
#elif DEF_16_GAP
    half16 v1 = vload16(id, nodeVar);
    v1 *= WeightsNode[id];
    atomiv_add(gnewNodeVar, v1.s0 + v1.s1 + v1.s2 + v1.s3 + v1.s4 + v1.s5 + v1.s6 + v1.s7 + v1.s8 + v1.s9 + v1.sa + v1.sb + v1.sc + v1.sd + v1.se + v1.sf);
    
#endif
    
    mem_fence(CLK_LOCAL_MEM_FENCE);
    if(get_local_id(0) == 0 ) {
        newNodeVar[ weightsOffset[id] ] = 1.0 / exp(gnewNodeVar);
    }
    
    
}



